#include "kernel_operator.h"

using namespace AscendC;

constexpr int32_t BUFFER_NUM = 1; 

template<typename typeT>
class KernelScatterElements {
public:
    __aicore__ inline KernelScatterElements() {}
    __aicore__ inline void Init(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, 
                                uint32_t varlength, uint32_t indiceslength, uint32_t updatelength, int32_t reduce, int32_t axis)
    {
        this->varlength = varlength;
        this->indiceslength = indiceslength;
        this->updatelength = updatelength;
        this->reduce = reduce;
        this->axis = axis;
        this->alignnum = 32 / sizeof(typeT);
        //this->rightPadding = this->alignnum - 1;

        varGm.SetGlobalBuffer((__gm__ typeT*)var, varlength + this->alignnum);
        indicesGm.SetGlobalBuffer((__gm__ int32_t*)indices, indiceslength + this->alignnum);
        updatesGm.SetGlobalBuffer((__gm__ typeT*)updates, updatelength + this->alignnum);      
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(inQueueUpdate, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(tmpBuffer1, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(QueueTmp1, this->alignnum * sizeof(float));
        pipe.InitBuffer(QueueTmp2, this->alignnum * sizeof(float));
        pipe.InitBuffer(tmp1, this->alignnum * sizeof(half));
        pipe.InitBuffer(tmp2, this->alignnum * sizeof(half));
        pipe.InitBuffer(tmp3, this->alignnum * sizeof(int16_t));
        pipe.InitBuffer(tmp6, this->alignnum * sizeof(int16_t));
        pipe.InitBuffer(tmp4, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(tmp5, this->alignnum * sizeof(typeT));
        // this->sumtmp = tmpBuffer1.Get<typeT>();
        // Duplicate(sumtmp, typeT(0), this->alignnum);
        // sumtmp.SetValue(0, typeT(1));

    }        
    __aicore__ inline void Process()
    {
        // if constexpr (std::is_same_v<typeT, int8_t> || std::is_same_v<typeT, signed char>)  {
        //     ComputeInt8();
        // } else {    
            for (uint32_t i = 0; i < this->indiceslength; i++) {
                int32_t index = indicesGm.GetValue(i);
                //typeT update = updatesGm.GetValue(i);
                CopyIn(index, i);
                Compute();
                CopyOut(index);
            }
        // }

    }
    __aicore__ inline void InitMultiDim(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, 
                                uint32_t varlength, uint32_t indiceslength, uint32_t updatelength, int32_t reduce, int32_t axis, int32_t dimnum,
                                int32_t sv[], int32_t si[], int32_t su[])
    {
        this->varlength = varlength;
        this->indiceslength = indiceslength;
        this->updatelength = updatelength;
        this->reduce = reduce;
        this->axis = axis;
        this->dimnum = dimnum;
        this->alignnum = 32 / sizeof(typeT);
        //this->rightPadding = this->alignnum - 1;
        for (int32_t i = 0; i < dimnum; ++i) {
            ((int32_t *)this->varshapeAccm)[i] = sv[i];
            ((int32_t *)this->indexshapeAccm)[i] = si[i];
            ((int32_t *)this->updateshapeAccm)[i] = su[i];
        }

        varGm.SetGlobalBuffer((__gm__ typeT*)var, varlength + this->alignnum);
        indicesGm.SetGlobalBuffer((__gm__ int32_t*)indices, indiceslength + this->alignnum);
        updatesGm.SetGlobalBuffer((__gm__ typeT*)updates, updatelength + this->alignnum);      
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(inQueueUpdate, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(tmpBuffer1, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(QueueTmp1, this->alignnum * sizeof(float));
        pipe.InitBuffer(QueueTmp2, this->alignnum * sizeof(float));
        pipe.InitBuffer(tmp1, this->alignnum * sizeof(half));
        pipe.InitBuffer(tmp2, this->alignnum * sizeof(half));
        pipe.InitBuffer(tmp3, this->alignnum * sizeof(int16_t));
        pipe.InitBuffer(tmp6, this->alignnum * sizeof(int16_t));
        pipe.InitBuffer(tmp4, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(tmp5, this->alignnum * sizeof(typeT));

    } 
    __aicore__ inline void ProcessMultiDim()
    {

        for (uint32_t i = 0; i < this->indiceslength; i++) {
            int32_t index = indicesGm.GetValue(i);
            GetVarPosNew(i, index);
            CopyIn(this->varpos, this->updatepos);
            Compute();
            CopyOut(this->varpos);
        }

    }  
    __aicore__ inline void InitNone(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, 
                                uint32_t varlength, uint32_t indiceslength, uint32_t updatelength)
    {
        this->varlength = varlength;
        this->indiceslength = indiceslength;
        this->updatelength = updatelength;
        this->alignnum = 32 / sizeof(typeT);

        varGm.SetGlobalBuffer((__gm__ typeT*)var, varlength + this->alignnum);
        indicesGm.SetGlobalBuffer((__gm__ int32_t*)indices, indiceslength + this->alignnum);
        updatesGm.SetGlobalBuffer((__gm__ typeT*)updates, updatelength + this->alignnum);      
        pipe.InitBuffer(inQUpdateBind, BUFFER_NUM, this->alignnum * sizeof(typeT));

    }    
    __aicore__ inline void ProcessNone()
    {

        for (uint32_t i = 0; i < this->indiceslength; i++) {
            int32_t index = indicesGm.GetValue(i);
            CopyInNone(i);
            CopyOutNone(index);
        }

    }  
    __aicore__ inline void InitNoneMultiDim(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, 
                                uint32_t varlength, uint32_t indiceslength, uint32_t updatelength, int32_t axis, int32_t dimnum,
                                int32_t sv[], int32_t si[], int32_t su[])
    {
        this->varlength = varlength;
        this->indiceslength = indiceslength;
        this->updatelength = updatelength;
        this->alignnum = 32 / sizeof(typeT);
        this->axis = axis;
        this->dimnum = dimnum;

        for (int32_t i = 0; i < dimnum; ++i) {
            ((int32_t *)this->varshapeAccm)[i] = sv[i];
            ((int32_t *)this->indexshapeAccm)[i] = si[i];
            ((int32_t *)this->updateshapeAccm)[i] = su[i];
        }

        varGm.SetGlobalBuffer((__gm__ typeT*)var, varlength + this->alignnum);
        indicesGm.SetGlobalBuffer((__gm__ int32_t*)indices, indiceslength + this->alignnum);
        updatesGm.SetGlobalBuffer((__gm__ typeT*)updates, updatelength + this->alignnum);      
        pipe.InitBuffer(inQUpdateBind, BUFFER_NUM, this->alignnum * sizeof(typeT));

    }  
    __aicore__ inline void ProcessNoneMultiDim()
    {

        for (uint32_t i = 0; i < this->indiceslength; i++) {
            int32_t index = indicesGm.GetValue(i);
            GetVarPosNew(i, index);
            // CopyInNone(this->updatepos);
            // CopyOutNone(this->varpos);
            typeT tmp = updatesGm.GetValue(this->updatepos);
            varGm.SetValue(this->varpos, tmp);

        }

    } 
    __aicore__ inline void InitNoneLastDim(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, 
                                uint32_t varlength, uint32_t indiceslength, uint32_t updatelength, 
                                int32_t varlastdimnum, int32_t indexlastdimnum, int32_t updatelastdimnum)
    {
        this->varlength = varlength;
        this->indiceslength = indiceslength;
        this->updatelength = updatelength;
        this->varlastdimnum = varlastdimnum;
        this->indexlastdimnum = indexlastdimnum;
        this->updatelastdimnum = updatelastdimnum;
        this->alignnum = 32 / sizeof(typeT);
        this->varlastdimnumalign = (varlastdimnum + this->alignnum - 1) / this->alignnum * this->alignnum;
        this->indexlastdimnumalign = (indexlastdimnum + this->alignnum - 1) / this->alignnum * this->alignnum;
        this->updatelastdimnumalign = (updatelastdimnum + this->alignnum - 1) / this->alignnum * this->alignnum;

        varGm.SetGlobalBuffer((__gm__ typeT*)var, varlength + this->alignnum);
        indicesGm.SetGlobalBuffer((__gm__ int32_t*)indices, indiceslength + this->alignnum);
        updatesGm.SetGlobalBuffer((__gm__ typeT*)updates, updatelength + this->alignnum);      
        pipe.InitBuffer(inQXBind, BUFFER_NUM, this->varlastdimnumalign * sizeof(typeT));
        pipe.InitBuffer(inQueueUpdate, BUFFER_NUM, this->updatelastdimnumalign * sizeof(typeT));
        pipe.InitBuffer(inQueueIndex, BUFFER_NUM, this->indexlastdimnumalign * sizeof(int32_t));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->varlastdimnumalign * sizeof(typeT));


    }  
    __aicore__ inline void ProcessNoneLastDim()
    {
        int32_t loopcount = this->indiceslength / this->indexlastdimnum;
        for (uint32_t i = 0; i < loopcount; i++) {
            CopyInNoneLastDim(i);
            ComputeNoneLastDim();
            CopyOutNoneLastDim(i);
        }

    } 

    __aicore__ inline void InitSclar(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, 
                                uint32_t varlength, uint32_t indiceslength, uint32_t updatelength, int32_t reduce, int32_t axis, int32_t dimnum,
                                int32_t sv[], int32_t si[], int32_t su[])
    {
        this->varlength = varlength;
        this->indiceslength = indiceslength;
        this->updatelength = updatelength;
        this->reduce = reduce;
        this->axis = axis;
        this->dimnum = dimnum;
        this->alignnum = 32 / sizeof(typeT);

        for (int32_t i = 0; i < dimnum; ++i) {
            ((int32_t *)this->varshapeAccm)[i] = sv[i];
            ((int32_t *)this->indexshapeAccm)[i] = si[i];
            ((int32_t *)this->updateshapeAccm)[i] = su[i];
        }

        varGm.SetGlobalBuffer((__gm__ typeT*)var, varlength + this->alignnum);
        indicesGm.SetGlobalBuffer((__gm__ int32_t*)indices, indiceslength + this->alignnum);
        updatesGm.SetGlobalBuffer((__gm__ typeT*)updates, 1);      
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(inQueueUpdate, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(tmpBuffer1, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(QueueTmp1, this->alignnum * sizeof(float));
        pipe.InitBuffer(QueueTmp2, this->alignnum * sizeof(float));
        pipe.InitBuffer(tmp1, this->alignnum * sizeof(half));
        pipe.InitBuffer(tmp2, this->alignnum * sizeof(half));
        pipe.InitBuffer(tmp3, this->alignnum * sizeof(int16_t));
        pipe.InitBuffer(tmp6, this->alignnum * sizeof(int16_t));
        pipe.InitBuffer(tmp4, this->alignnum * sizeof(typeT));
        pipe.InitBuffer(tmp5, this->alignnum * sizeof(typeT));

    } 
    __aicore__ inline void ProcessSclar()
    {

        for (uint32_t i = 0; i < this->indiceslength; i++) {
            int32_t index = indicesGm.GetValue(i);
            if(this->dimnum > 1) {
                GetVarPosSclar(i, index);
                typeT tmp = updatesGm.GetValue(0);
                if(this->reduce == 3) {
                    varGm.SetValue(this->varpos, tmp);
                } else {
                    ComputeSclar(i, this->varpos);
                }
            } else {
                typeT tmp = updatesGm.GetValue(0);
                if(this->reduce == 3) {
                    varGm.SetValue(index, tmp);
                } else {
                    ComputeSclar(i, index);
                }
            }
        }

    }  
private:   
 
    __aicore__ inline void GetVarPosNew(int32_t progrss, int32_t index) {
        this->varpos = 0;
        this->updatepos = 0;

        for(int32_t k = 0; k < this->dimnum; k++) {
            if(k == this->dimnum - 1) {
                if(k == this->axis) {
                    this->varpos += index;
                } else {
                    this->varpos += progrss;
                }
                this->updatepos += progrss;
            } else {
                int32_t dimrank = progrss / indexshapeAccm[k];
                progrss = progrss % indexshapeAccm[k];

                this->updatepos += dimrank * updateshapeAccm[k];

                if(k == this->axis) {
                    this->varpos += index * varshapeAccm[k];
                } else {
                    this->varpos += dimrank * varshapeAccm[k];
                }
            }
        }
    }
    __aicore__ inline void GetVarPosSclar(int32_t progrss, int32_t index) {
        this->varpos = 0;

        for(int32_t k = 0; k < this->dimnum; k++) {
            if(k == this->dimnum - 1) {
                if(k == this->axis) {
                    this->varpos += index;
                } else {
                    this->varpos += progrss;
                }

            } else {
                int32_t dimrank = progrss / indexshapeAccm[k];
                progrss = progrss % indexshapeAccm[k];

                if(k == this->axis) {
                    this->varpos += index * varshapeAccm[k];
                } else {
                    this->varpos += dimrank * varshapeAccm[k];
                }
            }
        }
    }    
    __aicore__ inline void GetVarPos(int32_t progrss, int32_t index) {
        this->varpos = 0;
        this->updatepos = 0;

        for(int32_t k = 0; k < this->dimnum; k++) {
            if(k == this->dimnum - 1) {
                if(k == this->axis) {
                    this->varpos += index;
                } else {
                    this->varpos += progrss;
                }
                this->updatepos += progrss;
            } else {
                int32_t sizetmp = 1;
                int32_t sizetmp2 = 1;
                int32_t sizetmp3 = 1;
                for(int32_t m = k + 1; m < this->dimnum; m++) {
                    sizetmp *= indexshape[m];
                    sizetmp2 *= updateshape[m];
                    sizetmp3 *= varshape[m];
                }
                int32_t dimrank = progrss / sizetmp;
                progrss = progrss % sizetmp;

                this->updatepos += dimrank * sizetmp2;

                if(k == this->axis) {
                    this->varpos += index * sizetmp3;
                } else {
                    this->varpos += dimrank * sizetmp3;
                }
            }
        }
    }
    __aicore__ inline void CopyIn(int32_t index, uint32_t progress) {
        LocalTensor<typeT> xLocal = inQueueX.AllocTensor<typeT>();
        DataCopy(xLocal, varGm[index], this->alignnum);
        inQueueX.EnQue(xLocal);

        LocalTensor<typeT> updateLocal = inQueueUpdate.AllocTensor<typeT>();
        //DataCopyExtParams copyParams{1, this->alignnum * sizeof(typeT), 0, 0, 0}; 
        DataCopyExtParams copyParams{1, 32, 0, 0, 0}; 
        DataCopyPadExtParams<typeT> padParams{false, 0, 0, 0};
        //DataCopyPadExtParams<typeT> padParams{true, 0, rightPadding, 0};
        DataCopyPad(updateLocal, updatesGm[progress], copyParams, padParams); 
        inQueueUpdate.EnQue(updateLocal);
    }
    __aicore__ inline void CopyOut(int32_t index)
    {
        // LocalTensor<typeT> yLocal = outQueueY.DeQue<typeT>();
        // DataCopy(varGm[index], yLocal, this->alignnum);   
        // outQueueY.FreeTensor(yLocal);

        LocalTensor<typeT> yLocal = outQueueY.DeQue<typeT>();
        DataCopyExtParams copyParams{1, 1 * sizeof(typeT), 0, 0, 0};
        DataCopyPad(varGm[index], yLocal, copyParams); 
        outQueueY.FreeTensor(yLocal);
    }
    __aicore__ inline void Compute() {
        LocalTensor<typeT> xLocal = inQueueX.DeQue<typeT>();
        LocalTensor<typeT> updateLocal = inQueueUpdate.DeQue<typeT>();
        LocalTensor<typeT> yLocal = outQueueY.AllocTensor<typeT>();
        if constexpr (std::is_same_v<typeT, int8_t> || std::is_same_v<typeT, signed char>)  {
            auto p1 = tmp1.Get<half>();
            auto p2 = tmp2.Get<half>();
            auto p3 = tmp3.Get<int16_t>();
            auto p6 = tmp6.Get<int16_t>();
            Cast(p1, xLocal, RoundMode::CAST_NONE, this->alignnum);
            Cast(p2, updateLocal, RoundMode::CAST_NONE, this->alignnum);
            if(this->reduce == 1) {
                Add(p2, p1, p2, 1);
                Cast(p1.ReinterpretCast<int16_t>(), p2, RoundMode::CAST_RINT, this->alignnum);
                ShiftLeft(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                ShiftRight(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                Cast(p2, p1.ReinterpretCast<int16_t>(), RoundMode::CAST_NONE, this->alignnum);
                Cast(yLocal, p2, RoundMode::CAST_NONE, this->alignnum);
            } else {
                /*
                Mul(p2, p1, p2, 1);
                Cast(p1.ReinterpretCast<int16_t>(), p2, RoundMode::CAST_RINT, this->alignnum);
                Duplicate(p2.ReinterpretCast<int16_t>(), int16_t(255), this->alignnum);
                And(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), p2.ReinterpretCast<int16_t>(), this->alignnum);
                // ShiftLeft(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                // ShiftRight(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                Cast(p2, p1.ReinterpretCast<int16_t>(), RoundMode::CAST_NONE, this->alignnum);
                Cast(yLocal, p2, RoundMode::CAST_NONE, this->alignnum);
                
                int8_t tmp = p1.ReinterpretCast<int16_t>().ReinterpretCast<int8_t>().GetValue(0);
                yLocal.SetValue(0, tmp);
                */
                Cast(p3, p1, RoundMode::CAST_RINT, this->alignnum);
                Cast(p6, p2, RoundMode::CAST_RINT, this->alignnum);               
                Mul(p3, p3, p6, 1);
                int8_t tmp = p3.ReinterpretCast<int8_t>().GetValue(0);
                yLocal.SetValue(0, tmp);

                // Mul(p2, p1, p2, 1);
                // Cast(p1.ReinterpretCast<int16_t>(), p2, RoundMode::CAST_RINT, this->alignnum);
                // ShiftLeft(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                // ShiftRight(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                // Cast(p2, p1.ReinterpretCast<int16_t>(), RoundMode::CAST_NONE, this->alignnum);
                // Cast(yLocal, p2, RoundMode::CAST_NONE, this->alignnum);
            }

            // Cast(p1.ReinterpretCast<int16_t>(), p2, RoundMode::CAST_RINT, this->alignnum);
            // ShiftLeft(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
            // ShiftRight(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
            // Cast(p2, p1.ReinterpretCast<int16_t>(), RoundMode::CAST_NONE, this->alignnum);
            // Cast(yLocal, p2, RoundMode::CAST_NONE, this->alignnum);
        } else if constexpr (std::is_same_v<typeT, half>)  {
            // LocalTensor<float> tmp1 = QueueTmp1.Get<float>();
            // LocalTensor<float> tmp2 = QueueTmp2.Get<float>();
            // Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->alignnum);
            // Cast(tmp2, updateLocal, RoundMode::CAST_NONE, this->alignnum);
            // Add(tmp1, tmp1, tmp2, 1);
            // Cast(yLocal, tmp1, RoundMode::CAST_TRUNC, this->alignnum);

            // typeT x1 = xLocal.GetValue(0);
            // typeT update = updateLocal.GetValue(0);
            // float x1fp32 = x1;
            // float updatefp32 = update;
            // x1fp32 += updatefp32;
            // typeT result = x1fp32;
            // yLocal.SetValue(0, result);
            if(this->reduce == 1) {
                Add(yLocal, xLocal, updateLocal, 1);
            } else {
                //Mul(yLocal, xLocal, updateLocal, 1);
                LocalTensor<float> tmp1 = QueueTmp1.Get<float>();
                LocalTensor<float> tmp2 = QueueTmp2.Get<float>();
                Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->alignnum);
                Cast(tmp2, updateLocal, RoundMode::CAST_NONE, this->alignnum);
                Mul(tmp1, tmp1, tmp2, 1);
                Cast(yLocal, tmp1, RoundMode::CAST_RINT, this->alignnum);
            }    
        } else {
            // LocalTensor<typeT> tmp = tmpBuffer1.Get<typeT>();
            // Duplicate(tmp, typeT(0), this->alignnum);
            // Duplicate(tmp, update, 1);
            // Add(xLocal, xLocal, tmp, this->alignnum);
            // inQueueX.EnQue(xLocal);
            if(this->reduce == 1) {
                Add(yLocal, xLocal, updateLocal, 1);
            } else {
                Mul(yLocal, xLocal, updateLocal, 1);
            }    
        }
        inQueueX.FreeTensor(xLocal);
        inQueueUpdate.FreeTensor(updateLocal);
        outQueueY.EnQue(yLocal);
    }

    __aicore__ inline void ComputeSclar(uint32_t progress, int32_t varpos) {
        if constexpr (std::is_same_v<typeT, int8_t> || std::is_same_v<typeT, signed char>)  {  
            
            int32_t tmpindex = indicesGm.GetValue(progress);
            int8_t tmpupdate = updatesGm.GetValue(0);
            int8_t tmpvar = varGm.GetValue(tmpindex);

            auto p1 = tmp1.Get<half>();
            auto p2 = tmp2.Get<half>();
            auto p3 = tmp3.Get<int16_t>();
            auto p6 = tmp6.Get<int16_t>();
            auto p4 = tmp4.Get<typeT>();
            auto p5 = tmp5.Get<typeT>();

            p4.SetValue(0, tmpupdate);
            p5.SetValue(0, tmpvar);
            Cast(p1, p4, RoundMode::CAST_NONE, this->alignnum);
            Cast(p2, p5, RoundMode::CAST_NONE, this->alignnum);
            if(this->reduce == 1) {
                Add(p2, p1, p2, 1);
                Cast(p1.ReinterpretCast<int16_t>(), p2, RoundMode::CAST_RINT, this->alignnum);
                ShiftLeft(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                ShiftRight(p1.ReinterpretCast<int16_t>(), p1.ReinterpretCast<int16_t>(), int16_t(8), this->alignnum);
                Cast(p2, p1.ReinterpretCast<int16_t>(), RoundMode::CAST_NONE, this->alignnum);
                Cast(p4, p2, RoundMode::CAST_NONE, this->alignnum);
                int8_t tmp = p4.GetValue(0);
                varGm.SetValue(varpos, tmp);
            } else {
                Cast(p3, p1, RoundMode::CAST_RINT, this->alignnum);
                Cast(p6, p2, RoundMode::CAST_RINT, this->alignnum);               
                Mul(p3, p3, p6, 1);
                int8_t tmp = p3.ReinterpretCast<int8_t>().GetValue(0);
                varGm.SetValue(tmpindex, tmp);
            }            
        } else {
            int32_t tmpindex = indicesGm.GetValue(progress);
            typeT tmpupdate = updatesGm.GetValue(0);
            typeT tmpvar = varGm.GetValue(tmpindex);

            auto p4 = tmp4.Get<typeT>();
            auto p5 = tmp5.Get<typeT>();
            p4.SetValue(0, tmpupdate);
            p5.SetValue(0, tmpvar);
            if(this->reduce == 1) {
                Add(p4, p4, p5, 1);
            } else {
                Mul(p4, p4, p5, 1);
            }
            typeT tmp = p4.GetValue(0);
            varGm.SetValue(varpos, tmp);
        }   
    }    
    __aicore__ inline void CopyInNone(uint32_t progress) {

        LocalTensor<typeT> updateLocal = inQUpdateBind.AllocTensor<typeT>();
        DataCopyExtParams copyParams{1, 32, 0, 0, 0}; 
        DataCopyPadExtParams<typeT> padParams{false, 0, 0, 0};
        DataCopyPad(updateLocal, updatesGm[progress], copyParams, padParams); 
        inQUpdateBind.EnQue(updateLocal);
    }  
    __aicore__ inline void CopyOutNone(int32_t index)
    {

        LocalTensor<typeT> updateLocal = inQUpdateBind.DeQue<typeT>();
        DataCopyExtParams copyParams{1, 1 * sizeof(typeT), 0, 0, 0};
        DataCopyPad(varGm[index], updateLocal, copyParams); 
        inQUpdateBind.FreeTensor(updateLocal);
    } 
    __aicore__ inline void CopyInNoneLastDim(uint32_t progress) {
        LocalTensor<typeT> xLocal = inQXBind.AllocTensor<typeT>();
        DataCopy(xLocal, varGm[progress * this->varlastdimnum], this->varlastdimnumalign);
        inQXBind.EnQue(xLocal);

        LocalTensor<typeT> updateLocal = inQueueUpdate.AllocTensor<typeT>();
        DataCopy(updateLocal, updatesGm[progress * this->updatelastdimnum], this->updatelastdimnumalign);
        inQueueUpdate.EnQue(updateLocal);

        LocalTensor<int32_t> indexLocal = inQueueIndex.AllocTensor<int32_t>();
        DataCopy(indexLocal, indicesGm[progress * this->indexlastdimnum], this->indexlastdimnumalign);
        inQueueIndex.EnQue(indexLocal);
    }  
    __aicore__ inline void ComputeNoneLastDim() {
        LocalTensor<typeT> xLocal = inQXBind.DeQue<typeT>();
        LocalTensor<typeT> updateLocal = inQueueUpdate.DeQue<typeT>();
        LocalTensor<int32_t> indexLocal = inQueueIndex.DeQue<int32_t>();
        //LocalTensor<typeT> yLocal = outQueueY.AllocTensor<typeT>();   

        // Muls(indexLocal, indexLocal, int32_t(sizeof(typeT)), this->indexlastdimnum);
        // Scatter(yLocal, updateLocal, indexLocal.ReinterpretCast<uint32_t>(), (uint32_t)0, this->indexlastdimnum);
        // if constexpr (std::is_same_v<typeT, int8_t> || std::is_same_v<typeT, signed char>)  {
        //     Muls(yLocal.ReinterpretCast<int16_t>(), xLocal.ReinterpretCast<int16_t>(), int16_t(1), this->varlastdimnumalign / 2);
        // } else {
        //     Muls(yLocal, xLocal, typeT(1), this->varlastdimnum);
        // }  
        //Muls(yLocal, xLocal, typeT(1), this->varlastdimnum);
        for(int32_t k = 0; k < this->indexlastdimnum; k++) {
            int32_t tmpindex = indexLocal.GetValue(k);
            typeT tmpupdate = updateLocal.GetValue(k);
            xLocal.SetValue(tmpindex, tmpupdate);
        }

        inQueueIndex.FreeTensor(indexLocal);
        inQueueUpdate.FreeTensor(updateLocal);        
        inQXBind.EnQue(xLocal);
        //outQueueY.EnQue(yLocal);
    } 
    __aicore__ inline void CopyOutNoneLastDim(uint32_t progress)
    {
        // LocalTensor<typeT> yLocal = outQueueY.DeQue<typeT>();
        // uint32_t blockLen = this->varlastdimnum * sizeof(typeT);
        // DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
        // DataCopyPad(varGm[progress * this->varlastdimnum], yLocal, copyParams); 
        // outQueueY.FreeTensor(yLocal);
        LocalTensor<typeT> xLocal = inQXBind.DeQue<typeT>();
        uint32_t blockLen = this->varlastdimnum * sizeof(typeT);
        DataCopyExtParams copyParams{1, blockLen, 0, 0, 0};
        DataCopyPad(varGm[progress * this->varlastdimnum], xLocal, copyParams); 
        inQXBind.FreeTensor(xLocal);

    }               
private:
    TPipe pipe;
    TBuf<QuePosition::VECCALC> tmpBuffer1; 
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2;
    TBuf<QuePosition::VECCALC> tmp1, tmp2, tmp3, tmp4, tmp5, tmp6;
    TQueBind<TPosition::VECIN, TPosition::VECOUT, BUFFER_NUM> inQUpdateBind;
    TQueBind<TPosition::VECIN, TPosition::VECOUT, BUFFER_NUM> inQXBind;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueUpdate;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIndex;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    GlobalTensor<typeT> varGm;
    GlobalTensor<int32_t> indicesGm;
    GlobalTensor<typeT> updatesGm;
    GlobalTensor<typeT> var_refGm;
    LocalTensor<typeT> sumtmp;
    uint32_t varlength;
    uint32_t indiceslength;
    uint32_t updatelength;
    int32_t reduce;
    int32_t axis;
    uint32_t alignnum;
    int32_t dimnum;
    int32_t varlastdimnum;
    int32_t indexlastdimnum;
    int32_t updatelastdimnum;
    int32_t varlastdimnumalign;
    int32_t indexlastdimnumalign;
    int32_t updatelastdimnumalign;
    int32_t varshape[64];
    int32_t indexshape[64];
    int32_t updateshape[64];
    int32_t varshapeAccm[64];
    int32_t indexshapeAccm[64];
    int32_t updateshapeAccm[64];      
    int32_t varpos;
    int32_t updatepos;
};

extern "C" __global__ __aicore__ void scatter_elements(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(t_data, tiling);
    KernelScatterElements<DTYPE_VAR> op;
    if (TILING_KEY_IS(100)) {
        op.Init(var, indices, updates,
                t_data.varlength, t_data.indiceslength, t_data.updatelength, t_data.reduce, t_data.axis);
        op.Process();
    } else if (TILING_KEY_IS(200)) {
        op.InitMultiDim(var, indices, updates,
                t_data.varlength, t_data.indiceslength, t_data.updatelength, t_data.reduce, t_data.axis, t_data.dimnum,
                t_data.varshapeAccm, t_data.indexshapeAccm, t_data.updateshapeAccm);
        op.ProcessMultiDim();
    } else if (TILING_KEY_IS(300)) {
        op.InitNone(var, indices, updates,
                    t_data.varlength, t_data.indiceslength, t_data.updatelength);
        op.ProcessNone();  
    } else if (TILING_KEY_IS(400)) {
        op.InitNoneMultiDim(var, indices, updates,
                    t_data.varlength, t_data.indiceslength, t_data.updatelength, t_data.axis, t_data.dimnum,
                    t_data.varshapeAccm, t_data.indexshapeAccm, t_data.updateshapeAccm);
        op.ProcessNoneMultiDim();     
    // } else if (TILING_KEY_IS(410)) {
    //     op.InitNoneLastDim(var, indices, updates,
    //                 t_data.varlength, t_data.indiceslength, t_data.updatelength,
    //                 t_data.varlastdimnum, t_data.indexlastdimnum, t_data.updatelastdimnum);
    //     op.ProcessNoneLastDim();  
    } else if (TILING_KEY_IS(500)) {
        op.InitSclar(var, indices, updates,
                    t_data.varlength, t_data.indiceslength, t_data.updatelength, t_data.reduce, t_data.axis, t_data.dimnum,
                    t_data.varshapeAccm, t_data.indexshapeAccm, t_data.updateshapeAccm);
        op.ProcessSclar();                      
    }   
}